#include "main.h"
#include <stdio.h>
#include <CL/cl.h>


char *c_kernel_source = {"\
bool is_equal(const class_T a, const class_T b){\
    return a == b;\
}\
__kernel __attribute__((reqd_work_group_size(blockdim, 1, 1)))\
void benchmark(class_T seed, global class_T *g_data){\
    const unsigned int blockSize = blockdim;\
	const int grid_size = blockSize * get_num_groups(0);\
	const int stride = grid_size;\
	const int idx = get_global_id(0);\
	const int big_stride = get_num_groups(0)*blockSize*ELEMENTS_PER_THREAD;\
	class_T tmps[ELEMENTS_PER_THREAD];\
	for(int k=0; k<FUSION_DEGREE; k++){	\
		for(int j=0; j<ELEMENTS_PER_THREAD; j++){\
			tmps[j] = g_data[idx+j*stride+k*big_stride];\
			for(int i=0; i<COMPUTE_ITERATIONS; i++){\
				tmps[j] = tmps[j]*tmps[j]+seed;\
			}\
		}\
		class_T sum = (class_T)0;\
		for(int j=0; j<ELEMENTS_PER_THREAD; j+=2)\
			sum += tmps[j]*tmps[j+1];\
        if(is_equal(sum, (class_T)-1))\
			g_data[idx+k*big_stride] = sum;\
	}\
}"};

struct Result SUCCESS(int runtime, int score){
	struct Result result = {runtime, score, 0};
	return result;
};

struct Result FAILURE(int state){
	struct Result result = {0, 0, state};
	return result;
};

struct Result DLL_EXPORT benchmark(){
    const char c_param_format_str[] = "-cl-std=CL1.2 -cl-mad-enable -Dclass_T=%s -Dblockdim=%d -DCOMPUTE_ITERATIONS=64 -DELEMENTS_PER_THREAD=%d -DFUSION_DEGREE=%d -DENABLE_HP";
    const int itersize = 10;
    const long size = 1024*1024*32;
    const int workgroupsize = 256;
    const int elements_per_wi = 8;
    const int fusion_degree = 4;
    const long compute_grid_size = size/elements_per_wi/fusion_degree;
    const size_t dimBlock[1] = {workgroupsize};
	const size_t dimGrid[1] = {(size_t)compute_grid_size};

    cl_uint deviceNum;
    cl_platform_id platform = NULL;
    cl_device_id device;
	cl_int status;

	status = clGetPlatformIDs(1, &platform, NULL);
	if (status != CL_SUCCESS){
        printf("Error: can not get platforms!\n");
        return FAILURE(status);
	}
    status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &deviceNum);
    if (status != CL_SUCCESS){
        printf("Error: can not find gpu device\n");
        return FAILURE(status);
    }
    if (deviceNum == 0){
        printf("WARNING: no gpu device available\n");
        return FAILURE(-1);
    }

    status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, deviceNum, &device, NULL);
    if (status != CL_SUCCESS){
        printf("Error: can not get default gpu device\n");
        return FAILURE(status);
    }

    cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, &status);
    if (status != CL_SUCCESS){
        printf("Error: can not create opencl context!\n");
        return FAILURE(status);
    }

    cl_command_queue commandQueue = clCreateCommandQueue(context, device, CL_QUEUE_PROFILING_ENABLE, &status);
	if (status != CL_SUCCESS){
        printf("Error: can not create opencl command queue!\n");
        return FAILURE(status);
    }

    cl_mem out_buffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, size*sizeof(double), NULL, &status);
    if (status != CL_SUCCESS){
        printf("Error: can not create out put buffer!\n");
        return FAILURE(status);
    }

    cl_int *mapped_data = (cl_int*)clEnqueueMapBuffer(commandQueue, out_buffer, CL_TRUE, CL_MAP_WRITE, 0, size*sizeof(double), 0, NULL, NULL, &status);
	if (status != CL_SUCCESS){
        printf("Error: can not create mapped data!\n");
        return FAILURE(status);
    }
	for(int i=0; i<size; i++){
        mapped_data[i] = 0;
	}
	clEnqueueUnmapMemObject(commandQueue, out_buffer, mapped_data, 0, NULL, NULL);



    double score_count = 0;
	double runtime_count = 0;
    //------------------------------- float ----------------------------------
    char build_params_float[256];
    sprintf(build_params_float, c_param_format_str, "float", workgroupsize, elements_per_wi, fusion_degree);
    cl_kernel kernels_float[itersize];
    for(int i=0; i < itersize; i++){
        cl_program program = clCreateProgramWithSource(context, 1, (char **)&c_kernel_source, NULL, &status);
        if (status != CL_SUCCESS){
            printf("Error: can not create program with opencl source!\n");
            return FAILURE(status);
        }
        status = clBuildProgram(program, 1, &device, build_params_float, NULL, NULL);
        if (status != CL_SUCCESS){
            printf("Error: can not build program with opencl source\n");
            return FAILURE(status);
        }

        cl_kernel kernel = clCreateKernel(program, "benchmark", &status);
        if (status != CL_SUCCESS){
            printf("Error: can not create opencl kernel\n");
            return FAILURE(status);
        }
        kernels_float[i] = kernel;
    }

    for(int i=0; i < itersize; i++){
        cl_kernel kernel = kernels_float[i];

        const cl_float seed_float = 1.0f + i;
        clSetKernelArg(kernel, 0, sizeof(cl_float), &seed_float);
        clSetKernelArg(kernel, 1, sizeof(cl_mem), &out_buffer);

        cl_event event;
        status = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, dimGrid, dimBlock, 0, NULL, &event);
        if (status != CL_SUCCESS){
            printf("Error: can not create opencl range kernel, error:%d!\n", status);
            return FAILURE(status);
        }
        clWaitForEvents(1, &event);

        cl_ulong time_start, time_end;
        clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
        clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);

        cl_program program;
        clGetKernelInfo(kernel, CL_KERNEL_PROGRAM, sizeof(program), &program, NULL);
        clReleaseEvent(event);
        clReleaseKernel(kernel);
        clReleaseProgram(program);

        double runtime = (time_end - time_start) / 1000000.0;
        double score = 20000.0 / runtime;

        score_count += score;
        runtime_count += runtime;
	}

    //------------------------------- double ----------------------------------
    char build_params_double[256];
    sprintf(build_params_double, c_param_format_str, "double", workgroupsize, elements_per_wi, fusion_degree);
    cl_kernel kernels_double[itersize];
    for(int i=0; i < itersize; i++){
        cl_program program = clCreateProgramWithSource(context, 1, (char **)&c_kernel_source, NULL, &status);
        if (status != CL_SUCCESS){
            printf("Error: can not create program with opencl source!\n");
            return FAILURE(status);
        }
        status = clBuildProgram(program, 1, &device, build_params_double, NULL, NULL);
        if (status != CL_SUCCESS){
            printf("Error: can not build program with opencl source\n");
            return FAILURE(status);
        }

        cl_kernel kernel = clCreateKernel(program, "benchmark", &status);
        if (status != CL_SUCCESS){
            printf("Error: can not create opencl kernel\n");
            return FAILURE(status);
        }
        kernels_double[i] = kernel;
    }

    for(int i=0; i < itersize; i++){
        cl_kernel kernel = kernels_double[i];

        const cl_double seed_double = 1.0 + i;
        clSetKernelArg(kernel, 0, sizeof(cl_double), &seed_double);
        clSetKernelArg(kernel, 1, sizeof(cl_mem), &out_buffer);

        cl_event event;
        status = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, dimGrid, dimBlock, 0, NULL, &event);
        if (status != CL_SUCCESS){
            printf("Error: can not create opencl range kernel, error:%d!\n", status);
            return FAILURE(status);
        }
        clWaitForEvents(1, &event);

        cl_ulong time_start, time_end;
        clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
        clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);

        cl_program program;
        clGetKernelInfo(kernel, CL_KERNEL_PROGRAM, sizeof(program), &program, NULL);
        clReleaseEvent(event);
        clReleaseKernel(kernel);
        clReleaseProgram(program);

        double runtime = (time_end - time_start) / 1000000.0;
        double score = 500000.0 / runtime;

        score_count += score;
        runtime_count += runtime;
	}

    //------------------------------- int ----------------------------------
    cl_kernel kernels_int[itersize];
    char build_params_int[256];
    sprintf(build_params_int, c_param_format_str, "int", workgroupsize, elements_per_wi, fusion_degree);
    for(int i=0; i < itersize; i++){
        cl_program program = clCreateProgramWithSource(context, 1, (char **)&c_kernel_source, NULL, &status);
        if (status != CL_SUCCESS){
            printf("Error: can not create program with opencl source!\n");
            return FAILURE(status);
        }
        status = clBuildProgram(program, 1, &device, build_params_int, NULL, NULL);
        if (status != CL_SUCCESS){
            printf("Error: can not build program with opencl source\n");
            return FAILURE(status);
        }

        cl_kernel kernel = clCreateKernel(program, "benchmark", &status);
        if (status != CL_SUCCESS){
            printf("Error: can not create opencl kernel\n");
            return FAILURE(status);
        }
        kernels_int[i] = kernel;
    }


	for(int i=0; i < itersize; i++){
        cl_kernel kernel = kernels_int[i];

        const cl_int seed_int = i;
        clSetKernelArg(kernel, 0, sizeof(cl_int), &seed_int);
        clSetKernelArg(kernel, 1, sizeof(cl_mem), &out_buffer);

        cl_event event;
        status = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, dimGrid, dimBlock, 0, NULL, &event);
        if (status != CL_SUCCESS){
            printf("Error: can not create opencl range kernel, error:%d!\n", status);
            return FAILURE(status);
        }
        clWaitForEvents(1, &event);

        cl_ulong time_start, time_end;
        clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_START, sizeof(time_start), &time_start, NULL);
        clGetEventProfilingInfo(event, CL_PROFILING_COMMAND_END, sizeof(time_end), &time_end, NULL);

        cl_program program;
        clGetKernelInfo(kernel, CL_KERNEL_PROGRAM, sizeof(program), &program, NULL);
        clReleaseEvent(event);
        clReleaseKernel(kernel);
        clReleaseProgram(program);

        double runtime = (time_end - time_start) / 1000000.0;
        double score = 10000.0 / runtime;

        score_count += score;
        runtime_count += runtime;
	}
	size_t  max_work_group_size;
	clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &max_work_group_size, NULL);
	double score_rate = score_count * max_work_group_size * 0.6 / 1024;

    printf("runtime                 : %d\n", (int) runtime_count);
	printf("score                   : %d\n", (int) score_rate / itersize);


	clFlush(commandQueue);
	clReleaseCommandQueue(commandQueue);
	clReleaseContext(context);

    return SUCCESS((int) runtime_count, (int) score_rate / itersize);
};